__kernel void matchingSimulatedModel(
	__global ushort* observeDepthMap,
	__global ushort* renderDepthMap, 
	__global ushort* observeSkin,
	const ushort Hmax,
	const ushort Hmin,
	__global ushort* sum_1,
	__global ushort* sum_2,
	__global ushort* sum_3)
{
	int num = get_global_id(0);
	
	ushort sub = abs(  (int)observeDepthMap[num] - (int)renderDepthMap[num]  )  ;
	if(sub > Hmax )
	{
		sum_1[num] = 4;
	}else{
		sum_1[num] = (sub*4/Hmax);
	}

	bool rMatch;
	if(sub < Hmin && renderDepthMap[num] < 0xffff){
		rMatch=true;
	}else{
		rMatch=false;
	}
	sum_2[num] = observeSkin[num] != 0xffff || rMatch;
	sum_3[num] = observeSkin[num] != 0xffff && rMatch;
}

__kernel void matchingObserve(
	__global ushort* observeDepthMap,
	__global ushort* renderDepthMap, 
	__global uchar* observeSkin,
	const ushort Hmax,
	const ushort Hmin,
	__global ushort* sum_1,
	__global ushort* sum_2,
	__global ushort* sum_3)
{
	int num = get_global_id(0);
	ushort sub = abs(  (int)observeDepthMap[num] - (int)renderDepthMap[num]  )  ;
	if(sub > Hmax )
	{
		sum_1[num] = 4;
	}else{
		sum_1[num] = (sub*4/Hmax);
	}

	bool rMatch;
	if(sub < Hmin && renderDepthMap[num] < 0xffff){
		rMatch=true;
	}else{
		rMatch=false;
	}
	sum_2[num] = (observeSkin[num] == 0xff) || rMatch;
	sum_3[num] = (observeSkin[num] == 0xff) && rMatch;
}

__kernel void reduction_vector(__global ushort8* data, 
      __local ushort8* partial_sums, __global ushort* output) {

   int lid = get_local_id(0);
   int group_size = get_local_size(0);

   partial_sums[lid] = data[get_global_id(0)];
   barrier(CLK_LOCAL_MEM_FENCE);

   for(int i = group_size/2; i>0; i >>= 1) {
      if(lid < i) {
         partial_sums[lid] += partial_sums[lid + i];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
   }

	if(lid == 0) {
      output[get_group_id(0)] = partial_sums[0].s0+partial_sums[0].s1
			+partial_sums[0].s2+partial_sums[0].s3+partial_sums[0].s4
			+partial_sums[0].s5+partial_sums[0].s6+partial_sums[0].s7;
	}
}